fix gpu out token counter update error#988
Conversation
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Summary of Changes
Hello @hiworldwzj, I'm Gemini Code Assist1! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!
This pull request addresses and fixes an error in the GPU-based output token counter update mechanism. It introduces a more robust and flexible way to update token counts on the GPU by allowing conditional updates via a mask, and integrates this improved counting logic throughout the inference pipeline to ensure accuracy.
Highlights
- GPU Token Counter Kernel Enhancement: The Triton kernel
_token_id_counter_update_kernelhas been updated to accept an optionalmask_ptrand aHAS_MASKconstexpr. This allows the atomic addition operation to be conditionally applied based on a per-element boolean mask, ensuring that only relevant tokens contribute to the counter, which is crucial for accurate counting in scenarios like speculative decoding. - Refactored Token Counter Management: The
ReqManagernow includes a new dedicated method,update_reqs_out_token_counter_gpu, which specifically handles GPU-based output token counter updates by calling the enhanced Triton kernel. This change cleanly separates GPU-based counting logic from CPU-based counting, improving code organization and maintainability. - Integration of Masked GPU Counting: The newly enhanced GPU token counter update mechanism, complete with masking support, has been integrated into various prefill and decode operations across different inference backends (e.g.,
chunked_prefillanddp_backend). This ensures that the GPU-side token counts are accurately maintained for all generated tokens, especially when only a subset of tokens (e.g., accepted tokens in speculative decoding) should be counted.
Using Gemini Code Assist
The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.
Invoking Gemini
You can request assistance from Gemini at any point in your pull request via creating an issue comment (i.e. comment on the pull request page) using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands.
| Feature | Command | Description |
|---|---|---|
| Code Review | /gemini review |
Performs a code review for the current pull request in its current state. |
| Pull Request Summary | /gemini summary |
Provides a summary of the current pull request in its current state. |
| Comment | @gemini-code-assist | Responds in comments when explicitly tagged, both in issue comments and review comments. |
| Help | /gemini help |
Displays a list of available commands. |
Customization
To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.
Limitations & Feedback
Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments or fill out our survey to provide feedback.
You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension.
Footnotes
-
Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩
There was a problem hiding this comment.
Code Review
This pull request fixes an issue with the GPU output token counter by adding the update logic in various model backends. The introduction of a mask for conditional updates is a good addition. The refactoring in req_manager.py improves code structure by separating GPU-specific logic. Overall, the changes are well-implemented and address the problem effectively. One suggestion improves code clarity and reduces duplication in the Triton kernel.
| if HAS_MASK: | ||
| mask = tl.load(mask_ptr + offs, mask=loc_mask, other=False) | ||
| tl.atomic_add( | ||
| req_to_out_token_id_counter_ptr + req_idx * counter_stride_m + token_ids * counter_stride_n, | ||
| 1, | ||
| mask=loc_mask & mask, | ||
| ) | ||
| else: | ||
| tl.atomic_add( | ||
| req_to_out_token_id_counter_ptr + req_idx * counter_stride_m + token_ids * counter_stride_n, | ||
| 1, | ||
| mask=loc_mask, | ||
| ) |
There was a problem hiding this comment.
To improve maintainability and reduce code duplication, this if/else block can be refactored. The tl.atomic_add call is repeated in both branches with only the mask parameter being different. Determine the mask to use first, and then make a single tl.atomic_add call.
final_mask = loc_mask
if HAS_MASK:
mask = tl.load(mask_ptr + offs, mask=loc_mask, other=False)
final_mask = final_mask & mask
tl.atomic_add(
req_to_out_token_id_counter_ptr + req_idx * counter_stride_m + token_ids * counter_stride_n,
1,
mask=final_mask,
)1 participant
No description provided.